; ModuleID = 'reduction.cpp'
source_filename = "reduction.cpp"
target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7"
target triple = "amdgcn-amd-amdhsa"

; Function Attrs: convergent mustprogress nofree norecurse nounwind memory(argmem: readwrite)
define protected amdgpu_kernel void @_Z23atomic_reduction_kernelPiS_i(ptr addrspace(1) nocapture noundef readonly %in.coerce, ptr addrspace(1) nocapture noundef %out.coerce, i32 noundef %ARRAYSIZE) local_unnamed_addr #0 {
entry:
  %call.i.i = tail call i64 @__ockl_get_group_id(i32 noundef 0) #3
  %conv.i.i = trunc i64 %call.i.i to i32
  %call.i.i14 = tail call i64 @__ockl_get_local_size(i32 noundef 0) #3
  %conv.i.i15 = trunc i64 %call.i.i14 to i32
  %mul = mul i32 %conv.i.i15, %conv.i.i
  %call.i.i16 = tail call i64 @__ockl_get_local_id(i32 noundef 0) #3
  %conv.i.i17 = trunc i64 %call.i.i16 to i32
  %add = add i32 %mul, %conv.i.i17
  %cmp22 = icmp slt i32 %add, %ARRAYSIZE
  br i1 %cmp22, label %for.body, label %for.cond.cleanup

for.cond.cleanup:                                 ; preds = %for.body, %entry
  %sum.0.lcssa = phi i32 [ 0, %entry ], [ %add5, %for.body ]
  %0 = atomicrmw add ptr addrspace(1) %out.coerce, i32 %sum.0.lcssa syncscope("agent-one-as") monotonic, align 4
  ret void

for.body:                                         ; preds = %entry, %for.body
  %i.024 = phi i32 [ %add9, %for.body ], [ %add, %entry ]
  %sum.023 = phi i32 [ %add5, %for.body ], [ 0, %entry ]
  %idxprom = sext i32 %i.024 to i64
  %arrayidx = getelementptr inbounds i32, ptr addrspace(1) %in.coerce, i64 %idxprom
  %1 = load i32, ptr addrspace(1) %arrayidx, align 4, !tbaa !6, !amdgpu.noclobber !10
  %add5 = add nsw i32 %1, %sum.023
  %call.i.i20 = tail call i64 @__ockl_get_num_groups(i32 noundef 0) #3
  %conv.i.i21 = trunc i64 %call.i.i20 to i32
  %mul8 = mul i32 %conv.i.i21, %conv.i.i15
  %add9 = add i32 %mul8, %i.024
  %cmp = icmp slt i32 %add9, %ARRAYSIZE
  br i1 %cmp, label %for.body, label %for.cond.cleanup, !llvm.loop !11
}

; Function Attrs: convergent mustprogress nofree norecurse nounwind memory(argmem: readwrite)
define protected amdgpu_kernel void @_Z24atomic_reduction_kernel2PiS_i(ptr addrspace(1) nocapture noundef readonly %in.coerce, ptr addrspace(1) nocapture noundef %out.coerce, i32 noundef %ARRAYSIZE) local_unnamed_addr #0 {
entry:
  %call.i.i = tail call i64 @__ockl_get_group_id(i32 noundef 0) #3
  %conv.i.i = trunc i64 %call.i.i to i32
  %call.i.i106 = tail call i64 @__ockl_get_local_size(i32 noundef 0) #3
  %conv.i.i107 = trunc i64 %call.i.i106 to i32
  %mul = mul i32 %conv.i.i107, %conv.i.i
  %call.i.i108 = tail call i64 @__ockl_get_local_id(i32 noundef 0) #3
  %conv.i.i109 = trunc i64 %call.i.i108 to i32
  %add = add i32 %mul, %conv.i.i109
  %mul5 = shl nsw i32 %add, 4
  %cmp114 = icmp slt i32 %mul5, %ARRAYSIZE
  br i1 %cmp114, label %for.body.lr.ph, label %for.cond.cleanup

for.body.lr.ph:                                   ; preds = %entry
  %mul69 = shl i32 %conv.i.i107, 4
  br label %for.body

for.cond.cleanup:                                 ; preds = %for.body, %entry
  %sum.0.lcssa = phi i32 [ 0, %entry ], [ %op.rdx117, %for.body ]
  %0 = atomicrmw add ptr addrspace(1) %out.coerce, i32 %sum.0.lcssa syncscope("agent-one-as") monotonic, align 4
  ret void

for.body:                                         ; preds = %for.body.lr.ph, %for.body
  %i.0116 = phi i32 [ %mul5, %for.body.lr.ph ], [ %add71, %for.body ]
  %sum.0115 = phi i32 [ 0, %for.body.lr.ph ], [ %op.rdx117, %for.body ]
  %idxprom = sext i32 %i.0116 to i64
  %arrayidx = getelementptr inbounds i32, ptr addrspace(1) %in.coerce, i64 %idxprom
  %1 = load <8 x i32>, ptr addrspace(1) %arrayidx, align 4, !tbaa !6
  %add34 = or i32 %i.0116, 8
  %idxprom35 = sext i32 %add34 to i64
  %arrayidx36 = getelementptr inbounds i32, ptr addrspace(1) %in.coerce, i64 %idxprom35
  %2 = load <8 x i32>, ptr addrspace(1) %arrayidx36, align 4, !tbaa !6
  %3 = tail call i32 @llvm.vector.reduce.add.v8i32(<8 x i32> %1)
  %4 = tail call i32 @llvm.vector.reduce.add.v8i32(<8 x i32> %2)
  %op.rdx = add i32 %3, %4
  %op.rdx117 = add i32 %op.rdx, %sum.0115
  %call.i.i112 = tail call i64 @__ockl_get_num_groups(i32 noundef 0) #3
  %conv.i.i113 = trunc i64 %call.i.i112 to i32
  %mul70 = mul i32 %mul69, %conv.i.i113
  %add71 = add i32 %mul70, %i.0116
  %cmp = icmp slt i32 %add71, %ARRAYSIZE
  br i1 %cmp, label %for.body, label %for.cond.cleanup, !llvm.loop !13
}

; Function Attrs: convergent mustprogress nofree norecurse nounwind memory(argmem: readwrite)
define protected amdgpu_kernel void @_Z24atomic_reduction_kernel3PiS_i(ptr addrspace(1) nocapture noundef readonly %in.coerce, ptr addrspace(1) nocapture noundef %out.coerce, i32 noundef %ARRAYSIZE) local_unnamed_addr #0 {
entry:
  %call.i.i = tail call i64 @__ockl_get_group_id(i32 noundef 0) #3
  %conv.i.i = trunc i64 %call.i.i to i32
  %call.i.i34 = tail call i64 @__ockl_get_local_size(i32 noundef 0) #3
  %conv.i.i35 = trunc i64 %call.i.i34 to i32
  %mul = mul i32 %conv.i.i35, %conv.i.i
  %call.i.i36 = tail call i64 @__ockl_get_local_id(i32 noundef 0) #3
  %conv.i.i37 = trunc i64 %call.i.i36 to i32
  %add = add i32 %mul, %conv.i.i37
  %mul5 = shl nsw i32 %add, 2
  %cmp42 = icmp slt i32 %mul5, %ARRAYSIZE
  br i1 %cmp42, label %for.body.lr.ph, label %for.cond.cleanup

for.body.lr.ph:                                   ; preds = %entry
  %mul21 = shl i32 %conv.i.i35, 2
  br label %for.body

for.cond.cleanup:                                 ; preds = %for.body, %entry
  %sum.0.lcssa = phi i32 [ 0, %entry ], [ %op.rdx, %for.body ]
  %0 = atomicrmw add ptr addrspace(1) %out.coerce, i32 %sum.0.lcssa syncscope("agent-one-as") monotonic, align 4
  ret void

for.body:                                         ; preds = %for.body.lr.ph, %for.body
  %i.044 = phi i32 [ %mul5, %for.body.lr.ph ], [ %add23, %for.body ]
  %sum.043 = phi i32 [ 0, %for.body.lr.ph ], [ %op.rdx, %for.body ]
  %idxprom = sext i32 %i.044 to i64
  %arrayidx = getelementptr inbounds i32, ptr addrspace(1) %in.coerce, i64 %idxprom
  %1 = load <4 x i32>, ptr addrspace(1) %arrayidx, align 4, !tbaa !6
  %2 = tail call i32 @llvm.vector.reduce.add.v4i32(<4 x i32> %1)
  %op.rdx = add i32 %2, %sum.043
  %call.i.i40 = tail call i64 @__ockl_get_num_groups(i32 noundef 0) #3
  %conv.i.i41 = trunc i64 %call.i.i40 to i32
  %mul22 = mul i32 %mul21, %conv.i.i41
  %add23 = add i32 %mul22, %i.044
  %cmp = icmp slt i32 %add23, %ARRAYSIZE
  br i1 %cmp, label %for.body, label %for.cond.cleanup, !llvm.loop !14
}

; Function Attrs: convergent mustprogress nofree nounwind willreturn memory(none)
declare hidden i64 @__ockl_get_group_id(i32 noundef) local_unnamed_addr #1

; Function Attrs: convergent mustprogress nofree nounwind willreturn memory(none)
declare hidden i64 @__ockl_get_local_size(i32 noundef) local_unnamed_addr #1

; Function Attrs: convergent mustprogress nofree nounwind willreturn memory(none)
declare hidden i64 @__ockl_get_local_id(i32 noundef) local_unnamed_addr #1

; Function Attrs: convergent mustprogress nofree nounwind willreturn memory(none)
declare hidden i64 @__ockl_get_num_groups(i32 noundef) local_unnamed_addr #1

; Function Attrs: nocallback nofree nosync nounwind willreturn memory(none)
declare i32 @llvm.vector.reduce.add.v8i32(<8 x i32>) #2

; Function Attrs: nocallback nofree nosync nounwind willreturn memory(none)
declare i32 @llvm.vector.reduce.add.v4i32(<4 x i32>) #2

attributes #0 = { convergent mustprogress nofree norecurse nounwind memory(argmem: readwrite) "amdgpu-flat-work-group-size"="1,1024" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx906" "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" "uniform-work-group-size"="true" }
attributes #1 = { convergent mustprogress nofree nounwind willreturn memory(none) "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx906" "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dot1-insts,+dot2-insts,+dot7-insts,+dpp,+gfx8-insts,+gfx9-insts,+s-memrealtime,+s-memtime-inst,+wavefrontsize64" }
attributes #2 = { nocallback nofree nosync nounwind willreturn memory(none) }
attributes #3 = { convergent nounwind willreturn memory(none) }

!llvm.module.flags = !{!0, !1, !2}
!opencl.ocl.version = !{!3}
!llvm.ident = !{!4, !5}

!0 = !{i32 1, !"amdgpu_code_object_version", i32 400}
!1 = !{i32 1, !"wchar_size", i32 4}
!2 = !{i32 8, !"PIC Level", i32 1}
!3 = !{i32 2, i32 0}
!4 = !{!"clang version 16.0.6 (https://github.com/llvm/llvm-project.git 7cbf1a2591520c2491aa35339f227775f4d3adf6)"}
!5 = !{!"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.2.4 24392 1e2c94795ee0d6ab8e2ff3035965a6b74e11b475)"}
!6 = !{!7, !7, i64 0}
!7 = !{!"int", !8, i64 0}
!8 = !{!"omnipotent char", !9, i64 0}
!9 = !{!"Simple C++ TBAA"}
!10 = !{}
!11 = distinct !{!11, !12}
!12 = !{!"llvm.loop.mustprogress"}
!13 = distinct !{!13, !12}
!14 = distinct !{!14, !12}
